
#include "add.h"



template<typename T>
__global__ void matrtix_add(T* A,T* B,T* C,int row,int col){
    int idx = blockIdx.x*blockDim.x + threadIdx.x;
    if(idx<row*col){
        int x = idx /col;
        int y = idx %col;
        C[x*col+y] = A[x*col+y] + B[x*col+y];
    }
}

template<>
__global__ void matrtix_add(half* A,half* B,half* C,int row,int col){
    int idx = blockIdx.x*blockDim.x + threadIdx.x;
    if(idx<row*col){
        int x = idx /col;
        int y = idx %col;
        C[x*col+y] = __hadd(A[x*col+y] , B[x*col+y]);
    }
}

template<typename T>
void launch_matrix_add(T* A,T* B,T* C ,int m,int k,int blocksize,int n){
    int blocknum = (n+blocksize-1)/blocksize;
    matrtix_add<T><<<blocknum,blocksize>>>(A,B,C,m,k);
}


template<typename T>
void CPU_Matrix_add(T* A,T* B,T* C ,int m,int k){
    for(int i=0;i<m;++i){
        for(int j=0;j<k;++j){
            C[i*k+j] = A[i*k+j] + B[i*k+j];
        }
    }
}

template<>
void CPU_Matrix_add(half* A,half* B,half* C ,int m,int k){
    for(int i=0;i<m;++i){
        for(int j=0;j<k;++j){
            C[i*k+j] = __float2half(__half2float(A[i*k+j]) + __half2float(B[i*k+j]));
        }
    }
}





// 显式实例化 float 和 half 类型,在 .cu 文件中，模板函数在编译阶段不会自动生成 float、half 等特定类型的实例，因此需要手动实例化。


template void launch_matrix_add<float>(float*, float*, float*, int, int, int, int);
template void launch_matrix_add<half>(half*, half*, half*, int, int, int, int);
template void launch_matrix_add<uint8_t>(uint8_t*, uint8_t*, uint8_t*, int, int, int, int);

template void CPU_Matrix_add<float>(float*, float*, float*, int, int);
template void CPU_Matrix_add<half>(half*, half*, half*, int, int);
template void CPU_Matrix_add<uint8_t>(uint8_t*, uint8_t*, uint8_t*, int, int);


